目標很簡單：先量再快。
這篇帶你一步一步用 Nsight Systems（看「整體流程／時間線」）和 Nsight Compute（看「單支 kernel 的瓶頸」），再搭配 程式內部量測、幾個經典優化手法與排錯清單。看完就能從「勉強能跑」，變成「知道為什麼慢、怎麼讓它快」。

0｜先建立正確的心法：先量測、再改動

• 不要一上來就改程式。先用工具找出最大的那塊時間（Hotspot），再對準它。
• 分清楚慢在哪：
  • Host↔Device 搬運慢（PCIe／Unified Memory 頻繁遷移）。
  • Kernel 計算慢（記憶體瓶頸、分歧、occupancy 低）。
  • 排程／啟動開銷（很多很小的 kernel、同步太多）。
• 80/20：常常 20% 的地方佔了 80% 的時間。

1｜整體鳥瞰：用 Nsight Systems 看時間線（CPU↔GPU）

你想知道：我的時間花在哪？拷貝多還是算多？kernel 能不能重疊？

1.1 安裝&啟動

• CUDA Toolkit 會附 CLI 工具：nsys。也有 GUI（Nsight Systems），安裝後可直接開啟 .nsys-rep 報告。

1.2 先用一次 CLI 快拍

nsys profile -t cuda,nvtx --stats=true -o report ./your_app
# 會產生 report.nsys-rep（GUI用）和摘要

1.3 看什麼？

• Time Chart / Timeline：一條龍看到 CPU thread、cudaMemcpy、各個 kernel 的開始/結束與是否重疊。
• CUDA API Summary：拷貝時間 vs kernel 時間比例。
• GPU Context/SM 利用率：是否長時間空轉。

1.4 立刻可行的動作

• 拷貝很肥？→ 合併多次小拷貝、改 page-locked（pinned）host memory、嘗試 streams 讓 HtoD 與 kernel 重疊。
• 看到一堆小 kernel？→ kernel fusion 或 CUDA Graphs 減少啟動開銷。
• CPU 等 GPU？→ 少用 cudaDeviceSynchronize()，改用 events 或 stream 回調，只在必要點同步。

範例圖：

2｜深入單顆 kernel：用 Nsight Compute 找瓶頸

你想知道：這支 kernel 到底卡在「記憶體」還是「算力」？warp 有沒有分歧？occupancy 夠不夠？

2.1 啟動（CLI 最快）

# 跑所有常用分析段落（第一次用這招最省事）
ncu --set full ./your_app

# 只針對名稱含 foo 的 kernel
ncu --set full --kernels ".*foo.*" ./your_app

也可用 GUI（Nsight Compute）直接打開 .ncu-rep，圖表更容易理解。

2.2 重點報表如何讀

• Summary：牆鐘時間、呼叫次數、平均耗時。

• Launch Statistics / Occupancy：

  • Achieved Occupancy（實際裝上 SM 的活躍比例）。太低 → 通常是暫存器或 shared memory 使用過多、block size 不理想。

• Memory Workload Analysis：

  • DRAM Throughput / L2 Hit Rate / Global Load Efficiency。
  • Global Memory Coalescing 若差，代表隔行/跳躍存取。

• Scheduler / Warp State：

  • Warp stall reasons（等待記憶體、同步、競爭…）。看「最常 Stall 的原因」，對症下藥。

• Speed of Light / Roofline（GUI）：

  • 告訴你是 memory-bound 還是 compute-bound。
  • 若落在記憶體區：先優化存取形狀與快取重用；
    若落在算力區：看是否能增加運算密度或用更快的數學近似/混合精度。

3｜程式內部量測（快速、好用、零學習成本）

3.1 CUDA Events 量某段 kernel 的時間

cudaEvent_t a,b; cudaEventCreate(&a); cudaEventCreate(&b);
cudaEventRecord(a);
myKernel<<<grid,block,shared,stream>>>(...);
cudaEventRecord(b); cudaEventSynchronize(b);
float ms=0; cudaEventElapsedTime(&ms,a,b);
printf("kernel took %.3f ms\n", ms);

優點：準確抓 GPU 執行時間。缺點：看不到細節，搭配 Nsight 才完整。

3.2 NVTX 標記區段（讓 Nsight Systems 的時間線有語意）

#include <nvToolsExt.h>
nvtxRangePushA("Stage: HtoD Copy");
// ... do copy ...
nvtxRangePop();

4｜實戰優化手冊（對照你的報表逐條做）

A）記憶體訪問（通常第一大宗）

1. Coalesced access（相鄰 thread 讀相鄰位址）：重排資料或索引，避免跳躍。

2. Shared Memory 做 tiling：重複使用的資料先搬到共享記憶體，少碰 DRAM。

   • 小心 bank conflict（同一 bank 被多 thread 同時存）。
   • 解法：在第二維加一點 padding（例如 tile[T][T+1]）。

3. Read-only broadcast 用 constant/texture cache：所有 thread 都讀同一數值時，constant memory 很省。

4. 避免無效搬運：把多次小 cudaMemcpy 合併成一次大拷貝；Host buffer 用 pinned memory。

5. Unified Memory 若造成大量 page fault：

   • 先用 cudaMemPrefetchAsync() 預取到裝置；或改回顯式 cudaMemcpy。

B）分歧（warp divergence）（Nsight Compute 會告訴你）

• 同一 warp 內如果 if 分兩路，就會排隊跑 → 用資料分群、分支轉算術（有時可用 ?: 或 max/min）、或把差異大的資料拆到不同 kernel。
• 迴圈跑次數差太多（像 fractal）：難避免，但可以讓「相似工作」落到同一 warp，減少嚴重分歧。

C）Occupancy（載入度）

• 用 cudaOccupancyMaxPotentialBlockSize() 或 Nsight 建議值找 block size sweet spot。

• 暫存器爆量會限縮 occupancy：

  • 觀察 -Xptxas -v 與 Nsight 的 Register Usage；
  • 適度拆函式、重用變數，或用 __restrict__ 幫編譯器別名分析。

• Shared memory 用太多也會卡 occupancy：檢討 tile 大小。

D）算力最佳化

• 混合精度：能接受的話，float 甚至 FP16/TF32（用對硬體和庫）可大幅提升吞吐。
• 使用 快速數學：-use_fast_math 或 __frcp_rn、__fsqrt_rn 等（需衡量精度）。
• 庫優先：矩陣/FFT/卷積 → 直接用 cuBLAS/cuFFT/cuDNN 幾乎一定比手寫快。

E）啟動開銷與管線化

• 很多小 kernel → kernel fusion（把連續的小步驟合併）；或使用 CUDA Graphs 批次提交。
• streams：HtoD／kernel／DtoH 分到不同 stream，讓 拷貝與計算重疊。
• 雙緩衝（double buffering）：一邊算一邊拷貝下一批資料。

F）多 GPU 與 NUMA

• 多 GPU：確認每張卡各自有工作，避免所有資料都擠在同一張。
• NUMA 主機：確保 CPU 執行緒綁到靠近該 GPU 的 NUMA node（進階主題，工具如 numactl）。

5｜以症狀找解方（速查表）

6｜一個「從 0 到優化」的實戰流程（照做就行）

1. Release 編譯：-O3 -lineinfo -Xptxas -v。

2. Nsight Systems：nsys profile -t cuda,nvtx --stats=true ./app

   • 認清「搬運 vs 計算」比例，先解決最大的那塊。

3. Nsight Compute（鎖定 1–2 支最花時間的 kernel）：ncu --set full --kernels "regex"

   • 判斷 memory-bound or compute-bound；看 coalescing、L2 命中、warp stall。

4. 小步優化：一次只改一件事（例如：把訪問改成連續、加 shared tile）。

5. 重量測：用相同輸入與環境，取多次平均，寫下 改善百分比。

6. 反覆：回到 2) 3) 再抓下一個瓶頸。

這就是專業團隊的常態工作流程。工具 + 方法論，沒有神奇捷徑。

延伸學習

• Nsight Compute 的 Roofline：把「每位元組可做多少 FLOPs（算術密度）」與硬體上限放在同一張圖，最直觀判斷你卡哪。
• Nsight Systems 的 GPU Trace：看每個 stream 的佔用與重疊品質。
• CUPTI / 自訂指標：進階使用情境，能做更細的線上監測。
• 庫與框架：善用 cuBLAS/cuFFT/cuDNN、CUTLASS、Thrust，少走重複造輪子的路。

一句話總結

先用 Nsight Systems 找「最大的慢點」→ 再用 Nsight Compute 拆解「單顆 kernel 的瓶頸」→ 按「記憶體、分歧、occupancy、算力、啟動」五大方向一條條對症下藥 → 量測、記錄、反覆。
當你會這一套，你已經具備把任何 CUDA 專案「從能跑，拉到會飛」的核心能力。